#ifndef RENDER_DEVICE_CU
#define RENDER_DEVICE_CU

#include "../Scene.hpp"
#include "../Renderer.hpp"
#include "../Object.hpp"
#include "../Triangle.hpp"

// // The main render function. This where we iterate over all pixels in the image,
// // generate primary rays and cast these rays into the scene. The content of the
// // framebuffer is saved to a file.

__global__ void kernel(Scene *scene, int spp, Vector3fDev *eye_pos,float scale,float imageAspectRatio,Vector3fDev* dev_framebuffer){
    int tid = blockIdx.x*blockDim.x+threadIdx.x;
    int offset = gridDim.x*blockDim.x;

    int total_count = scene->width*scene->height*spp;
    while(tid<total_count){
        int i = tid % scene->width;
        int j=(tid  / scene->width) % (scene->height);

        float x = (2 * (i + 0.5) / (float)scene->width - 1) *
                    imageAspectRatio * scale;
        float y = (1 - 2 * (j + 0.5) / (float)scene->height) * scale;
        Vector3fDev dir = normalize(Vector3fDev(-x, y, 1));

        dev_framebuffer[j * scene->width + i] = dev_framebuffer[j * scene->width + i] + scene->castRay(Ray(eye_pos, dir), 0,tid) / spp;
        tid += offset;
    }
}


void validGPUAvailable(){
    int count;
    cudaGetDeviceCount(&count);
    if(count==0){
        printf("no cuda gpu detected\n");
        exit(-1);
    }
    else{
        printf("cuda gpu available\n");
    }
    
}

void setCudaLimit(){
    size_t size_heap, size_stack;
    // 当出现内存错误时，适当调整此处参数
    cudaDeviceSetLimit(cudaLimitMallocHeapSize,10240000*8);
    cudaDeviceSetLimit(cudaLimitStackSize,32*1024);
    cudaDeviceGetLimit(&size_heap, cudaLimitMallocHeapSize);
    cudaDeviceGetLimit(&size_stack, cudaLimitStackSize);
    printf("Heap size found to be %d; Stack size found to be %d\n",(int)size_heap,(int)size_stack);
  }
  


void Renderer::Render_GPU(Scene *scene,int spp,Vector3fDev &eye_pos,Vector3fDev* framebuffer)
{
    validGPUAvailable();
    setCudaLimit();
    Scene *dev_scene;
    dev_scene=scene;
    // cudaMalloc((void**)&dev_scene, sizeof(Scene));
    // cudaMemcpy(dev_scene, scene, sizeof(Scene), cudaMemcpyHostToDevice);

    Vector3fDev *dev_framebuffer;
    cudaMalloc((void**)&dev_framebuffer, scene->height * scene->width * sizeof(Vector3fDev));
    cudaMemcpy(dev_framebuffer, framebuffer, scene->height * scene->width * sizeof(Vector3fDev), cudaMemcpyHostToDevice);

    Vector3fDev *dev_eye_pos;
    cudaMalloc((void**)&dev_eye_pos, sizeof(Vector3fDev));
    cudaMemcpy(dev_eye_pos, &eye_pos, sizeof(Vector3fDev), cudaMemcpyHostToDevice);


    float scale = tan_host(deg2rad_host(scene->fov * 0.5));
    float imageAspectRatio=scene->width / (float)scene->height;

    kernel<<<4096, 256>>>(dev_scene,spp, dev_eye_pos, scale, imageAspectRatio, dev_framebuffer);
    cudaDeviceSynchronize();
    cudaError_t err=cudaGetLastError();
    if(err!=cudaSuccess){
        printf("CUDA Error: %s\n", cudaGetErrorString(err));
        exit(-1);
    }
    cudaMemcpy(framebuffer, dev_framebuffer, scene->height * scene->width * sizeof(Vector3fDev), cudaMemcpyDeviceToHost);
    
    printf("Render_GPU done\n");
}

#endif